cuda : dynamic MMVQ nwarps for narrow matrices#20831
Open
JoursBleu wants to merge 1 commit intoggml-org:masterfrom
Open
cuda : dynamic MMVQ nwarps for narrow matrices#20831JoursBleu wants to merge 1 commit intoggml-org:masterfrom
JoursBleu wants to merge 1 commit intoggml-org:masterfrom
Conversation
github-actions
bot
added
Nvidia GPU
|
Thanks, JoursBleu. Can confirm this resolves the issue, also appears to be a slight improvement. original was 40.2 TPS now 41.14. |
Contributor
|
According to the llama.cpp AI usage policy:
Your posts very much read like they are machine-generated. Please clarify the extent to which language models were used. |
Contributor
Author
No, i didn't use AI to write the post. But I do write it in Chinese, and use Claude Opus to translate it to English. |
Contributor
|
Thank you for clarifying. The problem from our end as maintainers is that we get a lot of spam and the only feasible way for us to sift through it is to ban the use of language models altogether (since it is much easier to determine whether language models were used at all vs. whether someone carefully checked the language model outputs). |
JoursBleu
force-pushed
the
fix/mmvq-dynamic-nwarps-moe
branch
from
23bae61 to
afb6b78
Compare
Contributor
Author
|
Hi @IMbackK @JohannesGaessler, would please review when you have time, thanks! |
Contributor
|
I have not forgotten about this PR but there happen to be other, concurrent PRs that tough the same code. One of them in particular is relevant for getting good scaling with tensor parallelism so I'm prioritizing it over this PR. |
Contributor
|
Can you rebase on top of master? I'll then review this PR so that it van be merged. |
JoursBleu
force-pushed
the
fix/mmvq-dynamic-nwarps-moe
branch
from
afb6b78 to
71a7fb1
Compare
Contributor
Author
|
@JohannesGaessler rebased |
Contributor
JohannesGaessler
left a comment
JohannesGaessler
left a comment
There was a problem hiding this comment.
Please change the name of the function to calc_max_nwarps to better reflect how it is being used. The PR seems to be correct based on static code analysis but I will need to check it for regressions.
JoursBleu
force-pushed
the
fix/mmvq-dynamic-nwarps-moe
branch
2 times, most recently
from
a19a90b to
8cd2a84
Compare
JoursBleu
force-pushed
the
fix/mmvq-dynamic-nwarps-moe
branch
from
8cd2a84 to
40dc8a1
Compare
|
Can confirm the updated version still works, without this 36.66 TPS vs 41.22. |
Contributor
Author
|
hi @JohannesGaessler,
|
Contributor
|
When I tested the performance: Details
I am still seeing performance regressions for hardware that would be supposedly unaffected. And for RDNA4 the speedup is not consistent. So this PR should not be merged. |
Contributor
Actually, it seems that someone else was using the machine simultaneously so that was probably the reason. Still, I don't think the additional complexity is worthwhile if the speedup is inconsistent. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Fix MMVQ TG regression on MoE models from #19478.
#19478 increased
nwarpsto 8 on RDNA3/RDNA4 to better utilize memory bandwidth for bs=1 decode. However,nwarps=8assumes wide weight matrices. MoE expert FFN layers are narrow (512–2048 cols), so most warps have no work but still pay__syncthreads()and shared-memory reduction overhead, causing a net TG regression. This patch dynamically clampsnwarpsbased on the actual matrix width to avoid this.R9700 (gfx1201, RDNA4), ROCm 7.2
MoE (regression fixed):
Dense (no regression):
Full whitelist sweep — llama-2-7b, 1x R9700, tg512, r=5:
W7900 (gfx1100, RDNA3), ROCm 7.1
MoE (regression fixed):
Full whitelist sweep — llama-2-7b, 1x W7900, tg512, r=5:
Note: PR description translated with AI assistance.